extern "C" __device__ 
void draw_line(uchar4* buffer, int width, int height, int x0, int y0, int x1, int y1, uchar4 color) {
    int dx = abs(x1 - x0);
    int dy = -abs(y1 - y0);
    int sx = x0 < x1 ? 1 : -1;
    int sy = y0 < y1 ? 1 : -1;
    int err = dx + dy;
    
    while (true) {
        if (x0 >= 0 && x0 < width && y0 >= 0 && y0 < height) {
            int idx = y0 * width + x0;
            buffer[idx] = color;
        }
        if (x0 == x1 && y0 == y1) break;
        int e2 = 2 * err;
        if (e2 >= dy) {
            err += dy;
            x0 += sx;
        }
        if (e2 <= dx) {
            err += dx;
            y0 += sy;
        }
    }
}

extern "C" __global__ 
void render_model_kernel(float* model_verts, unsigned int* model_face_vrts, char* frame_buffer, int face_num, int width, int height)
{
    int face_idx = blockIdx.x * blockDim.x + threadIdx.x; 
    if (face_idx < face_num) 
    {
        unsigned int v1_idx = model_face_vrts[face_idx * 3];
        unsigned int v2_idx = model_face_vrts[face_idx * 3 + 1];
        unsigned int v3_idx = model_face_vrts[face_idx * 3 + 2];

        float v1_x = model_verts[v1_idx * 4];
        float v1_y = model_verts[v1_idx * 4 + 1];

        float v2_x = model_verts[v2_idx * 4];
        float v2_y = model_verts[v2_idx * 4 + 1];

        float v3_x = model_verts[v3_idx * 4];
        float v3_y = model_verts[v3_idx * 4 + 1];

        int v1_x_int = (v1_x + 1.0) * width / 2.0;
        int v1_y_int = (v1_y + 1.0) * height / 2.0;

        int v2_x_int = (v2_x + 1.0) * width / 2.0;
        int v2_y_int = (v2_y + 1.0) * height / 2.0;

        int v3_x_int = (v3_x + 1.0) * width / 2.0;
        int v3_y_int = (v3_y + 1.0) * height / 2.0;

        bool legal = v1_x_int >= 0 && v1_x_int < width && v1_y_int >=0 && v1_y_int < height;
        legal = legal && v2_x_int >= 0 && v2_x_int < width && v2_y_int >=0 && v2_y_int < height;
        legal = legal && v3_x_int >= 0 && v3_x_int < width && v3_y_int >=0 && v3_y_int < height;

        if (legal)
        {
            uchar4 color;
            color.x = 255;
            color.y = 0;
            color.z = 0;
            color.w = 255;

            draw_line((uchar4*)frame_buffer, width, height, v1_x_int, v1_y_int, v2_x_int, v2_y_int, color);
            draw_line((uchar4*)frame_buffer, width, height, v2_x_int, v2_y_int, v3_x_int, v3_y_int, color);
            draw_line((uchar4*)frame_buffer, width, height, v3_x_int, v3_y_int, v1_x_int, v1_y_int, color);

            int v1_fb_idx = v1_x_int + v1_y_int * width;
            int v2_fb_idx = v2_x_int + v2_y_int * width;
            int v3_fb_idx = v3_x_int + v3_y_int * width;

            color.x = 255;
            color.y = 255;
            color.z = 255;
            color.w = 255;

            uchar4* fb = (uchar4 *)frame_buffer;
            fb[v1_fb_idx] = color;
            fb[v2_fb_idx] = color;
            fb[v3_fb_idx] = color;
        }
    }
}